[NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel #958

carlushuang · 2021-05-28T14:48:19Z

This is PR of asm igemm for NHWC layout. Currently this is for fwd/bwd, fp32/fp16.

NHWC fwd fp32/fp16 kernel and solver ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC
NHWC bwd fp32/fp16 kernel and solver ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC
make sure working properly in both Heuristic mode and tuning mode
add ctest for testing

tested these cases: perf_data
Most of the files in this PR is generated by iGEMMgen, with a single .s file for a single kernel. Apart from this, only several cpp files have changes. Below is list of non-asm files that have changes or added

src/CMakeLists.txt
src/conv/invokers/impl_gemm_dynamic.cpp
src/include/miopen/conv/asm_implicit_gemm.hpp
src/include/miopen/conv/invokers/impl_gemm_dynamic.hpp
src/include/miopen/solver.hpp
src/mlo_dir_conv.cpp
src/solver.cpp
src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp
src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp
src/solver/conv_asm_implicit_gemm_gtc_perf_config.cpp
test/CMakeLists.txt

…_bwd

ghost · 2021-05-29T11:13:11Z

Congratulations 🎉. DeepCode analyzed your code in 2.149 seconds and we found no issues. Enjoy a moment of no bugs ☀️.

👉 View analysis in DeepCode’s Dashboard | Configure the bot

👉 The DeepCode service and API will be deprecated in August, 2021. Here is the information how to migrate. Thank you for using DeepCode 🙏 ❤️ !

If you are using our plugins, you might be interested in their successors: Snyk's JetBrains plugin and Snyk's VS Code plugin.

atamazov · 2021-05-31T21:56:29Z

@carlushuang Could you please add urgency and value labels. Thanks.

carlushuang · 2021-06-04T01:52:18Z

@atamazov Hi please help review this PR, now this PR have passed CI for once and framework team need this feature to do NHWC e2e test

jerryyin · 2021-06-04T13:58:41Z

I have little idea of how to interpret/validate assembly kernels. Abstain from reviewing this PR, but thanks for your contribution.

…side lambda

…_bwd

asroy

LGTM. I've only reviewed the host side code.

src/conv/invokers/impl_gemm_dynamic.cpp

…_bwd

junliume · 2021-06-18T01:41:55Z

[2021-06-17T09:09:23.503Z] Backward weights convolution: ConvAsmBwdWrW1x1
[2021-06-17T09:09:23.503Z] Input tensor: 64, 128, 28, 28
[2021-06-17T09:09:23.503Z] Weights tensor: 512, 128, 1, 1
[2021-06-17T09:09:23.503Z] Output tensor: 64, 512, 28, 28
[2021-06-17T09:09:23.503Z] Filter: conv2d, miopenConvolution, miopenPaddingDefault, {0, 0}, {1, 1}, {1, 1},
[2021-06-17T09:09:23.503Z] ../bin/test_conv2d --float --cmode conv --pmode default --group-count 1 --input 64, 160, 73, 73 --weights 64, 160, 1, 1 --pads_strides_dilations 0 0 1 1 1 1 --trans_output_pads 0 0 --in_layout NCHW --fil_layout NCHW --out_layout NCHW
[2021-06-17T09:09:23.503Z] FAILED: 0.000103983
[2021-06-17T09:09:23.503Z] Max diff: 738
[2021-06-17T09:09:23.503Z] Mismatch at 940430: -146 != -427

carlushuang · 2021-06-18T08:08:05Z

CI all passed
@junliume the fail about ConvAsmBwdWrW1x1 is not introduced by this PR. (This is direct algorithm), and after I re-start that test, all passed

atamazov · 2021-06-19T13:19:37Z

@junliume Thanks for the info from CI! Please note that correct quotation from the log is:

[2021-06-17T09:09:23.503Z] ../bin/test_conv2d --float --cmode conv --pmode default --group-count 1 --input 64, 160, 73, 73 --weights 64, 160, 1, 1 --pads_strides_dilations 0 0 1 1 1 1 --trans_output_pads 0 0 --in_layout NCHW --fil_layout NCHW --out_layout NCHW 
[2021-06-17T09:09:23.503Z] FAILED: 0.000103983
[2021-06-17T09:09:23.503Z] Max diff: 738
[2021-06-17T09:09:23.503Z] Mismatch at 940430: -146 != -427
[2021-06-17T09:09:23.503Z] Forward convolution: ConvOclDirectFwd1x1
[2021-06-17T09:09:23.503Z] Input tensor: 64, 160, 73, 73
[2021-06-17T09:09:23.503Z] Weights tensor: 64, 160, 1, 1
[2021-06-17T09:09:23.503Z] Output tensor: 64, 64, 73, 73
[2021-06-17T09:09:23.503Z] Filter: conv2d, miopenConvolution, miopenPaddingDefault, {0, 0}, {1, 1}, {1, 1},

and faulty Solver is ConvOclDirectFwd1x1.

atamazov

Blocked by #991

test/CMakeLists.txt

…_bwd

atamazov

None objections!

atamazov · 2021-06-23T23:29:29Z

@carlushuang ...but I still need an answer for this: #958 (comment)

# RESOLVED Conflicts: # src/solver.cpp

atamazov · 2021-06-23T23:38:39Z

CI run 30 has passed all the tests.

* add NHWC fwd/bwd fp32/fp16 kernel * fix several bug in Herustic and Tuning * add missing kernel and update config list * remove useless config lists * fix fwd fp32 not valid config * fix isValid * fix bwd fp16 not proper IsApplicable * add ctest for nhwc asm kernels * reorg NextLinear * parse in opArgs as mutable, reserve place for in/wei/out pointer outside lambda * Assign OpKernelArg() to opArgs vector, set proper ctest flag

carlushuang marked this pull request as draft May 28, 2021 14:48

carlushuang changed the title ~~[NHWC] fwd/bwd fp32/fp16 kernel~~ [NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel May 28, 2021

carlushuang added 2 commits May 28, 2021 22:58

add NHWC fwd/bwd fp32/fp16 kernel

ebdca83

Merge remote-tracking branch 'origin/develop' into asm_igemm_nhwc_fwd…

40ee16d

…_bwd

carlushuang added 6 commits May 29, 2021 19:24

fix several bug in Herustic and Tuning

f7e026f

add missing kernel and update config list

0b388cb

remove useless config lists

80abf01

fix fwd fp32 not valid config

b165d56

fix isValid

eea7469

fix bwd fp16 not proper IsApplicable

6e4e404

carlushuang marked this pull request as ready for review May 31, 2021 07:37

carlushuang added 4 commits May 31, 2021 15:50

add ctest for nhwc asm kernels

54b1cd2

fix clang tidy

fb0137b

further fix tidy

7f80f4f

reorg NextLinear

bbdbccd

carlushuang requested review from atamazov, jerryyin and junliume May 31, 2021 11:51

carlushuang added urgency_normal value_middle value_high and removed value_middle labels Jun 1, 2021

This comment has been minimized.

Sign in to view

carlushuang added TESTING_CI_PASSED urgency_high and removed urgency_normal labels Jun 3, 2021

carlushuang added 3 commits June 15, 2021 14:59

parse in opArgs as mutable, reserve place for inp/wei/out pointer out…

ff21946

…side lambda

Merge remote-tracking branch 'origin/develop' into asm_igemm_nhwc_fwd…

9bb27d2

…_bwd

fix tidy by assign const pointer properly

c3e3e0b

asroy previously approved these changes Jun 15, 2021

View reviewed changes

src/conv/invokers/impl_gemm_dynamic.cpp Outdated Show resolved Hide resolved

src/conv/invokers/impl_gemm_dynamic.cpp Outdated Show resolved Hide resolved

carlushuang dismissed asroy’s stale review via e085b49 June 16, 2021 02:56

remove useless code

e085b49

carlushuang requested a review from atamazov June 16, 2021 10:31

atamazov suggested changes Jun 16, 2021

View reviewed changes

src/conv/invokers/impl_gemm_dynamic.cpp Outdated Show resolved Hide resolved

src/conv/invokers/impl_gemm_dynamic.cpp Outdated Show resolved Hide resolved

carlushuang added 2 commits June 16, 2021 21:51

Merge remote-tracking branch 'origin/develop' into asm_igemm_nhwc_fwd…

66db338

…_bwd

Assign OpKernelArg() to opArgs vector, set proper ctest flag

a6ffe3b

junliume previously approved these changes Jun 18, 2021

View reviewed changes

carlushuang requested a review from atamazov June 19, 2021 03:13

atamazov mentioned this pull request Jun 19, 2021

ConvOclDirectFwd1x1: correctness error with "-y 1 -x 1 -n 64 -c 160 -H 73 -W 73 -k 64 -p 0 -q 0" #988

Closed

atamazov suggested changes Jun 20, 2021

View reviewed changes

test/CMakeLists.txt Outdated Show resolved Hide resolved

atamazov mentioned this pull request Jun 21, 2021

[TESTS][FP16][WORKAROUND] Fix test_conv_igemm_dynamic_xdlops* and test_regression_half_vega. W/A for #995 and #996 #991

Merged

Merge remote-tracking branch 'origin/develop' into asm_igemm_nhwc_fwd…

21bcc15

…_bwd

carlushuang dismissed junliume’s stale review via 21bcc15 June 23, 2021 08:46

atamazov previously approved these changes Jun 23, 2021

View reviewed changes

Merge branch 'develop' into asm_igemm_nhwc_fwd_bwd

8875ffc

# RESOLVED Conflicts: # src/solver.cpp

atamazov dismissed their stale review via 8875ffc June 23, 2021 23:37

atamazov approved these changes Jun 23, 2021

View reviewed changes

atamazov merged commit 6c0b289 into develop Jun 23, 2021

atamazov deleted the asm_igemm_nhwc_fwd_bwd branch June 23, 2021 23:40

carlushuang mentioned this pull request Jun 24, 2021

[ASM IGEMM] refactor asm igemm solver to have less copy inside invoker #1005

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel #958

[NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel #958

carlushuang commented May 28, 2021 •

edited

Loading

ghost commented May 29, 2021 •

edited by ghost

Loading

atamazov commented May 31, 2021

This comment has been minimized.

carlushuang commented Jun 4, 2021

jerryyin commented Jun 4, 2021

asroy left a comment

junliume commented Jun 18, 2021

carlushuang commented Jun 18, 2021

atamazov commented Jun 19, 2021

atamazov left a comment

atamazov left a comment

atamazov commented Jun 23, 2021

atamazov commented Jun 23, 2021

[NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel #958

[NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel #958

Conversation

carlushuang commented May 28, 2021 • edited Loading

ghost commented May 29, 2021 • edited by ghost Loading

👉 View analysis in DeepCode’s Dashboard | Configure the bot

👉 The DeepCode service and API will be deprecated in August, 2021. Here is the information how to migrate. Thank you for using DeepCode 🙏 ❤️ !

atamazov commented May 31, 2021

This comment has been minimized.

carlushuang commented Jun 4, 2021

jerryyin commented Jun 4, 2021

asroy left a comment

Choose a reason for hiding this comment

junliume commented Jun 18, 2021

carlushuang commented Jun 18, 2021

atamazov commented Jun 19, 2021

atamazov left a comment

Choose a reason for hiding this comment

atamazov left a comment

Choose a reason for hiding this comment

atamazov commented Jun 23, 2021

atamazov commented Jun 23, 2021

carlushuang commented May 28, 2021 •

edited

Loading

ghost commented May 29, 2021 •

edited by ghost

Loading